///**
// * @file main.cpp
// *
// * Copyright (C) 2024. Huawei Technologies Co., Ltd. All rights reserved.
// *
// * This program is distributed in the hope that it will be useful,
// * but WITHOUT ANY WARRANTY; without even the implied warranty of
// * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
// */
//#include "data_utils.h"
//#include "leakyrelu_custom_tilling.h"
//#ifndef ASCENDC_CPU_DEBUG
//#include "acl/acl.h"
//extern void leakyrelu_custom_do(uint32_t blockDim, void *stream, uint8_t *src, uint8_t *dst, float negativeSlope, LeakyReluCustomTilingData tiling);
//#else
//#include "tikicpulib.h"
//extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR src, GM_ADDR dst, float negativeSlope, LeakyReluCustomTilingData tiling);
//#endif
//extern void GenerateTilingData(uint8_t* tilingBuf);
//constexpr int32_t CORE_NUM = 8;
//
//int32_t main(int32_t argc, char *argv[])
//{
//    size_t inputByteSize = 8 * 2048 * sizeof(float);
//    size_t outputByteSize = 8 * 2048 * sizeof(float);
//    float negativeSlope = 0.1;
//    uint8_t *tiling = nullptr;
//    size_t tilingSize = sizeof(LeakyReluCustomTilingData);
//
//#ifdef ASCENDC_CPU_DEBUG
//    tiling = (uint8_t *)AscendC::GmAlloc(tilingSize);
//
//    uint8_t *src = (uint8_t *)AscendC::GmAlloc(inputByteSize);
//    uint8_t *dst = (uint8_t *)AscendC::GmAlloc(outputByteSize);
//
//    GenerateTilingData(tiling);
//
//    ReadFile("./input/input_src.bin", inputByteSize, src, inputByteSize);
//
//    AscendC::SetKernelMode(KernelMode::AIV_MODE);
//    ICPU_RUN_KF(leakyrelu_custom, CORE_NUM, src, dst, negativeSlope,
//                *reinterpret_cast<LeakyReluCustomTilingData *>(tiling));  // CPU Debug模式下核函数调用
//    WriteFile("./output/output_dst.bin", dst, outputByteSize);
//
//    AscendC::GmFree((void *)src);
//    AscendC::GmFree((void *)dst);
//    AscendC::GmFree((void *)tiling);
//#else
//    CHECK_ACL(aclInit(nullptr));
//    int32_t deviceId = 0;
//    CHECK_ACL(aclrtSetDevice(deviceId));
//    aclrtStream stream = nullptr;
//    CHECK_ACL(aclrtCreateStream(&stream));
//
//    uint8_t *srcHost, *dstHost;
//    uint8_t *srcDevice, *dstDevice;
//
//    CHECK_ACL(aclrtMallocHost((void **)(&srcHost), inputByteSize));
//    CHECK_ACL(aclrtMallocHost((void **)(&dstHost), outputByteSize));
//    CHECK_ACL(aclrtMallocHost((void **)(&tiling), tilingSize));
//    CHECK_ACL(aclrtMalloc((void **)&srcDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
//    CHECK_ACL(aclrtMalloc((void **)&dstDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
//
//    GenerateTilingData(tiling);
//
//    ReadFile("./input/input_src.bin", inputByteSize, srcHost, inputByteSize);
//
//    CHECK_ACL(aclrtMemcpy(srcDevice, inputByteSize, srcHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
//
////    ACLRT_LAUNCH_KERNEL(leakyrelu_custom)(CORE_NUM, stream, srcDevice, dstDevice, negativeSlope, *reinterpret_cast<LeakyReluCustomTilingData *>(tiling));
//    leakyrelu_custom_do(CORE_NUM, stream, srcDevice, dstDevice, negativeSlope, *reinterpret_cast<LeakyReluCustomTilingData>(tiling));
//    CHECK_ACL(aclrtSynchronizeStream(stream));
//
//    CHECK_ACL(aclrtMemcpy(dstHost, outputByteSize, dstDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));
//    WriteFile("./output/output_dst.bin", dstHost, outputByteSize);
//
//    CHECK_ACL(aclrtFree(srcDevice));
//    CHECK_ACL(aclrtFree(dstDevice));
//    CHECK_ACL(aclrtFreeHost(srcHost));
//    CHECK_ACL(aclrtFreeHost(dstHost));
//
//    CHECK_ACL(aclrtDestroyStream(stream));
//    CHECK_ACL(aclrtResetDevice(deviceId));
//    CHECK_ACL(aclFinalize());
//#endif
//    return 0;
//}

/**
 * @file main.cpp
 *
 * Copyright (C) 2025. Huawei Technologies Co., Ltd. All rights reserved.
 *
 * This program is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
 */
#include "leakyrelu_custom_tiling.h"
#include "data_utils.h"
#ifndef ASCENDC_CPU_DEBUG
#include "acl/acl.h"
#include "aclrtlaunch_leakyrelu_custom.h"
#include "tiling/platform/platform_ascendc.h"
#else
#include "tikicpulib.h"
extern "C" __global__ __aicore__ void leakyrelu_custom(GM_ADDR x, GM_ADDR z, float negativeSlope, LeakyReluCustomTilingData tilling);
#endif
extern void GenerateTilingData(uint8_t* tilingBuf, uint32_t blockDim);


int32_t main(int32_t argc, char *argv[])
{
    constexpr uint32_t BLOCK_DIM = 8;
    constexpr uint32_t DATA_TYPE_SIZE[] = {2, 2, 4, 1, 2, 4};
    uint8_t *tiling = nullptr;
    size_t tilingSize = 17 * sizeof(uint32_t);
    float negativeSlope = 0.1;


#ifdef ASCENDC_CPU_DEBUG
    tiling = (uint8_t *)AscendC::GmAlloc(tilingSize);
    ReadFile("./input/input_tiling.bin", tilingSize, tiling, tilingSize);
#else
    CHECK_ACL(aclInit(nullptr));
    int32_t deviceId = 0;
    CHECK_ACL(aclrtSetDevice(deviceId));
    aclrtStream stream = nullptr;
    CHECK_ACL(aclrtCreateStream(&stream));

    uint8_t *xHost, *zHost;
    uint8_t *xDevice, *zDevice;

    CHECK_ACL(aclrtMallocHost((void **)(&tiling), tilingSize));
    ReadFile("./input/input_tiling.bin", tilingSize, tiling, tilingSize);
#endif
    GenerateTilingData(tiling, BLOCK_DIM);
    uint32_t dataTypeSize = DATA_TYPE_SIZE[reinterpret_cast<LeakyReluCustomTilingData *>(tiling)->dataType];
    size_t inputByteSize = reinterpret_cast<LeakyReluCustomTilingData *>(tiling)->totalLength * dataTypeSize;
    size_t outputByteSize = reinterpret_cast<LeakyReluCustomTilingData *>(tiling)->totalLength * dataTypeSize;

#ifdef ASCENDC_CPU_DEBUG
    constexpr uint32_t DATA_BLOCK_BYTE = 32;
    uint8_t *x = (uint8_t *)AscendC::GmAlloc((inputByteSize + DATA_BLOCK_BYTE - 1) / DATA_BLOCK_BYTE * DATA_BLOCK_BYTE);
    uint8_t *z = (uint8_t *)AscendC::GmAlloc((inputByteSize + DATA_BLOCK_BYTE - 1) / DATA_BLOCK_BYTE * DATA_BLOCK_BYTE);

    ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);

    AscendC::SetKernelMode(KernelMode::AIV_MODE);

    ICPU_RUN_KF(leakyrelu_custom, BLOCK_DIM, x, z, negativeSlope,
                *reinterpret_cast<LeakyReluCustomTilingData *>(tiling));

    WriteFile("./output/output_z.bin", z, outputByteSize);

    AscendC::GmFree((void *)x);
    AscendC::GmFree((void *)z);
    AscendC::GmFree((void *)tiling);
#else
    CHECK_ACL(aclrtMallocHost((void **)(&xHost), inputByteSize));
    CHECK_ACL(aclrtMallocHost((void **)(&zHost), outputByteSize));
    CHECK_ACL(aclrtMalloc((void **)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
    CHECK_ACL(aclrtMalloc((void **)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

    ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);

    CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));

    ACLRT_LAUNCH_KERNEL(leakyrelu_custom)(BLOCK_DIM, stream, xDevice, zDevice, negativeSlope,
        reinterpret_cast<LeakyReluCustomTilingData *>(tiling));
    CHECK_ACL(aclrtSynchronizeStream(stream));

    CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));
    WriteFile("./output/output_z.bin", zHost, outputByteSize);

    CHECK_ACL(aclrtFree(xDevice));
    CHECK_ACL(aclrtFree(zDevice));
    CHECK_ACL(aclrtFreeHost(xHost));
    CHECK_ACL(aclrtFreeHost(zHost));
    CHECK_ACL(aclrtFreeHost(tiling));

    CHECK_ACL(aclrtDestroyStream(stream));
    CHECK_ACL(aclrtResetDevice(deviceId));
    CHECK_ACL(aclFinalize());
#endif
    return 0;
}
